[MIOpen] Port LRN kernels from OpenCL to HIP#4887
Merged
BradPepersAMD merged 5 commits intoROCm:developfrom Mar 12, 2026
Merged
[MIOpen] Port LRN kernels from OpenCL to HIP#4887BradPepersAMD merged 5 commits intoROCm:developfrom
BradPepersAMD merged 5 commits intoROCm:developfrom
Conversation
1 task
Contributor
There was a problem hiding this comment.
Pull request overview
Ports MIOpen’s LRN kernels from OpenCL-C to HIP source kernels and updates the host-side construction/invocation accordingly.
Changes:
- Replaced
MIOpenLRNFwd.cl/MIOpenLRNBwd.clwith HIP equivalents (.cpp) and updated CMake sources. - Updated LRN kernel construction to emit new compile-time defines and reference the new kernel filenames.
- Removed the unused
alphakernel argument from forward invocation and added a couple of HIP math helpers.
Reviewed changes
Copilot reviewed 9 out of 9 changed files in this pull request and generated 6 comments.
Show a summary per file
| File | Description |
|---|---|
| projects/miopen/src/ocl/mloNorm.cpp | Switches kernel filenames to .cpp, renames/reshapes compile-time defines, and updates forward/backward construction functions. |
| projects/miopen/src/ocl/lrn_ocl.cpp | Removes unused forward kernel parameter and updates a comment to be backend-agnostic. |
| projects/miopen/src/kernels/hip_math_ops.hpp | Adds device helpers for division/modulo style operations used by HIP kernels. |
| projects/miopen/src/kernels/MIOpenLRNFwd.cpp | Adds HIP implementation of LRN forward kernels (within-channel + across-channels). |
| projects/miopen/src/kernels/MIOpenLRNFwd.cl | Removes OpenCL forward kernel source. |
| projects/miopen/src/kernels/MIOpenLRNBwd.cpp | Adds HIP implementation of LRN backward kernels. |
| projects/miopen/src/kernels/MIOpenLRNBwd.cl | Removes OpenCL backward kernel source. |
| projects/miopen/src/include/miopen/mlo_internal.hpp | Updates LRN construct method signatures to void. |
| projects/miopen/src/CMakeLists.txt | Replaces .cl kernel sources with new .cpp kernel sources. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
BradPepersAMD
approved these changes
Mar 11, 2026
Removed unused kernel parameter from forward kernels. Tested by running `test_lrn` on MI210.
Requires also lifting an if condition out of the loop when specific conditions are met
* Fix typos in comments * Change C-style casts for reinterpret cast * comment about 24-bit range
EwanC
force-pushed
the
EwanC/miope_lrn_kernel_port
branch
from
737e954 to
eb13453
Compare
jovanau
pushed a commit
to jovanau/rocm-libraries
that referenced
this pull request
Mar 19, 2026
## Motivation Convert LRN MIOpen kernel from OpenCL-C to HIP ## Technical Details * Removed unused kernel parameter from forward kernels. * No equivalent of `mad24` from OpenCL-C in HIP, so used `__mul24` then add. * Had to use manual loop peeling and constexprs to replicate the loop unrolling behavior from the OpenCL compiler. ## Test Plan Tested by running `test_lrn` and doing a performance comparison. ## Test Result All tests pass on MI210 & MI300 and performance comparison to be sent internally. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
johannes-graner
pushed a commit
that referenced
this pull request
Mar 20, 2026
## Motivation Convert LRN MIOpen kernel from OpenCL-C to HIP ## Technical Details * Removed unused kernel parameter from forward kernels. * No equivalent of `mad24` from OpenCL-C in HIP, so used `__mul24` then add. * Had to use manual loop peeling and constexprs to replicate the loop unrolling behavior from the OpenCL compiler. ## Test Plan Tested by running `test_lrn` and doing a performance comparison. ## Test Result All tests pass on MI210 & MI300 and performance comparison to be sent internally. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Convert LRN MIOpen kernel from OpenCL-C to HIP
Technical Details
mad24from OpenCL-C in HIP, so used__mul24then add.Test Plan
Tested by running
test_lrnand doing a performance comparison.Test Result
All tests pass on MI210 & MI300 and performance comparison to be sent internally.
Submission Checklist